//! This is a silly hack to make Eclipse's syntax highlighting work with OpenCL kernels
#ifdef __CDT_PARSER__
#define __kernel
#define __global
#define __local
#endif


#define EDGE_THRESHOLD 5

#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable




__kernel void noop() {
        int i = 0;
        i +=5;
}

/*
*Fucntion will build 3 histograms from the values in the RGB channels of the image
*The histograms contain the frequency of occurence of each color for each channel
*
*/
__kernel void build_color_histograms (__read_only image2d_t inputImage,	
				  	sampler_t sampler,
									  __global int *r_color_histogram,
									  __global int *g_color_histogram,
									  __global int *b_color_histogram,
									  int single_image_width,
									  int width
									  ) 
{

	//since each operation is independent of all operations and includes
	//a single read that will not be used by other kernels, we will
	//not worry about local memory for this kernel
	
	//get the coord that this image will read
	int gid = get_global_id(0);
	int2 imageCoord = (int2)( gid - (gid/width) * width ,gid/width);
	
	uint4 pixelValue =  read_imageui(inputImage, sampler, imageCoord);
	
	
	//this array must have been initialized to 0 outside of the function!
	//we will ignore the fourth channel of the image (alpha)
	
	//some complicated math to find out what image we are currently in and then
	//which histogram index we should update
	int row_number = gid/width;
	int column_number = gid - row_number * width;
	int image_number = (gid - row_number * width) / single_image_width; 
	int offset = image_number * 256;
	
	
	//s2 is red
	atom_add(&r_color_histogram[pixelValue.s2 +offset ],1);
	
	//s1 is green
	atom_add(&g_color_histogram[pixelValue.s1 + offset],1);
	
	//s0 is blue
	atom_add(&b_color_histogram[pixelValue.s0 + offset],1);

}



//we can/must run three instances of this, one for each channel since the channels will operate independently.
//we cannot run more than one of these on the same channel for an image
//This function may fail (produce incorrect output, not crash) if some histograms indices have VERY very large numbers of pixels.
//Because of things like jpeg noise, this is NOT likely to occur, even in images that have very similar colors


__kernel void get_colors_from_histograms (
									  __global int *single_channel_histogram,
									  __global int *output_array,
									  int pixels_per_color,
									  int num_colors
									  ) {

	int gid = get_global_id(0);
    	int current_sum = 0;
	int num_pixels = 0;
	int hist_index = 0;
	int offset = gid * 256;
	int output_index = num_colors * gid;
	
	for (hist_index = 0 ; hist_index<256 ; hist_index++){
		int pixels_at_index = single_channel_histogram[hist_index+offset]; 
		num_pixels += pixels_at_index;
		current_sum += pixels_at_index * hist_index;
		
		if (num_pixels >= pixels_per_color){
			//find the average of those pixels and write it to the output array
			output_array[output_index++] = current_sum / num_pixels;			
			num_pixels = 0;
			current_sum = 0;
		}
	}
	
	//now write the final point(s) if it/they was not written in the loop above.  
	//In most case this will only write the LAST index, but in some cases it may write more
	for(; output_index < (num_colors * (gid+1)); output_index++){
		if (num_pixels > 0)
			output_array[output_index]=current_sum/num_pixels;
		else output_array[output_index]=output_array[output_index-1];
	}
	//the function is now finished.  The output_array contains all of the new colors discovered by averahing pixels in a group
}


/*
*This function will colorize a single pixel based on the color that it most closely matches in the new colors array.  We will colorize all three channels of the pixel
*/
__kernel void build_image_from_new_colors(	__read_only image2d_t inputImage,
											__write_only image2d_t outputImage,
											sampler_t sampler,
											__global int *r_colors_array,
											__global int *g_colors_array,
											__global int *b_colors_array,
											__global bool * edges,
											int num_colors,
											int single_image_width,
											int width
											) 
	{

		//we should probably cooperatively build up the RGB array into local memory.  We will do that in a later version of the kernel
		
		

		//first, get the current value of the pixel
		int gid = get_global_id(0);
		int2 imageCoord = (int2)( gid - (gid/width) * width ,gid/width);
		
		int im_number = (gid - (gid/width) * width) / single_image_width;

		
		//first check to see if the pixel should be black (an edge) or not.  If the pixel is an edge, then we don't need
		//to do any of the calculations below
		uint4 pixelValue = read_imageui(inputImage, sampler, imageCoord);

		if (edges[imageCoord.x + width * imageCoord.y] == true){
			write_imageui(outputImage, imageCoord, (uint4)(0, 0, 0, 255));
		}
		else{
			//now we want to find which color in the colors array is closest
			//as the baseline, we set the FIRST color in the r, g, and b arrays as closest
			int best_red = r_colors_array[im_number*num_colors];
			int best_green = g_colors_array[im_number*num_colors];
			int best_blue = b_colors_array[im_number*num_colors];
		
			int i = 1;
			uint temp_red, temp_green, temp_blue;
			uint  dist_red, dist_green, dist_blue;
			for(; i < num_colors ; i++){
				temp_red = abs_diff(r_colors_array[im_number*num_colors+i], (int)pixelValue.s0);   	
		        	temp_green = abs_diff(g_colors_array[im_number*num_colors+i], (int)pixelValue.s1);
				temp_blue = abs_diff(b_colors_array[im_number*num_colors+i], (int)pixelValue.s2);
			
				dist_red = abs_diff(best_red, (int)pixelValue.s0);
				dist_green = abs_diff(best_green, (int)pixelValue.s1);
				dist_blue = abs_diff(best_blue, (int)pixelValue.s2);
				//now check to see if any of these values are closer than the previous ones
				if (temp_red < dist_red){
					best_red = r_colors_array[im_number*num_colors+i];
				}
				if (temp_green < dist_green){
					best_green = g_colors_array[im_number*num_colors+i];
				}			
				if (temp_blue < dist_blue){
					best_blue = b_colors_array[im_number*num_colors+i];
				}
			
			}
		
			//finally, write the best colors back to the image
			//is the alpha layer 0 or 1 for no transparency???  For now use pixelValue.s3
		
			write_imageui(outputImage, imageCoord, (uint4)(best_red, best_green, best_blue, 255));
		}
		
	}




// Gaussian filter of image
//Slightly adapted from http://code.google.com/p/opencl-book-samples/
__kernel void gaussian_filter(__read_only image2d_t srcImg,
                              __write_only image2d_t dstImg,
                              sampler_t sampler,
                              int width, int height)
{
    // Gaussian Kernel is:
    // 1  2  1
    // 2  4  2
    // 1  2  1
    int kernelWeights[9] = {  1, 2, 1,
                               2, 4, 2,
                               1, 2, 1 };

    int2 startImageCoord = (int2) (get_global_id(0) - 1, get_global_id(1) - 1);
    int2 endImageCoord   = (int2) (get_global_id(0) + 1, get_global_id(1) + 1);
    int2 outImageCoord = (int2) (get_global_id(0), get_global_id(1));

    if (outImageCoord.x < width && outImageCoord.y < height)
    {
        int weight = 0;
        uint4 outColor = (uint4)(0, 0, 0, 0);
        for( int y = startImageCoord.y; y <= endImageCoord.y; y++)
        {
            for( int x = startImageCoord.x; x <= endImageCoord.x; x++)
            {
                outColor += (read_imageui(srcImg, sampler, (int2)(x, y)) * kernelWeights[weight] / 16);
                weight += 1;
            }
        }

        // Write the output value to image
        write_imageui(dstImg, outImageCoord, outColor);
    }
}


//Edge detection kernel
//Slightly adapted version of adapted gaussian kernel originally found at http://code.google.com/p/opencl-book-samples/
__kernel void sobel(__read_only image2d_t srcImg,
					 __global bool * line,
					 sampler_t sampler,
					 int width, int height)
{
	float v_kernelWeights[9] = {-1.0f, 0.0f, 1.0f,
				    -2.0f, 0.0f, 2.0f,
			            -1.0f, 0.0f, 1.0f};
	float h_kernelWeights[9] = {-1.0f, -2.0f, -1.0f,
		        	    0.0f, 0.0f, 0.0f,
			            1.0f, 2.0f, 1.0f};

	int2 startImageCoord = (int2) (get_global_id(0)-1, get_global_id(1)-1);

	int2 endImageCoord = (int2) (get_global_id(0)+1, get_global_id(1)+1);

	int2 outImageCoord = (int2) (get_global_id(0), get_global_id(1));

	float4 outColor1 = (float4)(0.0f, 0.0f, 0.0f, 0.0f);
	if (outImageCoord.x < width && outImageCoord.y < height)
	{
		int weight = 0;
		for(int y=startImageCoord.y; y <= endImageCoord.y; y++)
		{
			for(int x= startImageCoord.x; x<= endImageCoord.x; x++)
			{
				 outColor1 +=
					(convert_float4(read_imageui(srcImg, sampler, (int2)(x,y))) *
					v_kernelWeights[weight] / 16.0);
				weight += 1;
			}
		}
	}






	float4 outColor2 = (float4)(0.0f, 0.0f, 0.0f, 0.0f);
	if (outImageCoord.x < width && outImageCoord.y < height)
	{
		int weight = 0;
		for(int y=startImageCoord.y; y <= endImageCoord.y; y++)
		{
			for(int x= startImageCoord.x; x<= endImageCoord.x; x++)
			{
				outColor2 +=
					(convert_float4(read_imageui(srcImg, sampler, (int2)(x,y))) *
					h_kernelWeights[weight] / 16.0);
				weight += 1;
			}
		}
	}
	// write the output value ot the image
	float4 finalColor = sqrt((outColor1*outColor1)+(outColor2*outColor2));

	if (fabs(finalColor.x) > EDGE_THRESHOLD &&
		fabs(finalColor.y) > EDGE_THRESHOLD &&
		fabs(finalColor.z) > EDGE_THRESHOLD)
	{
		//since there is no easy way to pass a 2D array, we have passed in a 1D
		//array and we calculate the address like this
		line[outImageCoord.x + outImageCoord.y * width] = true;
	}
		
	//instead of doing the below step in the kernel, we will do it in the host.
	//at first kernel call, all array indices will be initialized to zero.  Since
	//global memory is slow, it is done in the host instead of the kernel
	/*else{
		line[outImageCoord.x,outImageCoord.y] = false;
	}
	*/
	
}
